extern "C" {

typedef long IndexType;
typedef long DvmhLoopRef;

void *dvmh_get_natural_base_(long *deviceRef, long dvmDesc[]);
void dvmh_fill_header_(long *deviceRef, void *base, long dvmDesc[], long dvmhDesc[]);
long loop_do_cuda(DvmhLoopRef *InDvmhLoop, dim3 *OutBlocks, dim3 *OutThreads, cudaStream_t *OutStream, IndexType **InOutBlocks);
void loop_register_red_cuda(DvmhLoopRef *InDvmhLoop, int InRedNum, void **ArrayPtr, void **LocPtr);
void loop_red_init(DvmhLoopRef *InDvmhLoop, int InRedNum, void *arrayPtr, void *locPtr);

void init_cuda_() {
}

void loop__1_cuda_(long *cur_dev, long *loop_ref, long a[], long b[])
{
    long da[5], db[5];
    dim3 blocks, threads;
    cudaStream_t stream;
    IndexType *blocks_info;
    void *a_base, *b_base;
// Get 'natural' bases
    a_base = dvmh_get_natural_base_(cur_dev, a);
    b_base = dvmh_get_natural_base_(cur_dev, b);
// Fill 'device' headers
    dvmh_fill_header_(cur_dev, a_base, a, da);
    dvmh_fill_header_(cur_dev, b_base, b, db);
// GPU execution
    while (loop_do_cuda(loop_ref, &blocks, &threads, &stream, &blocks_info) != 0)
        loop__1<<<blocks, threads, 0, stream>>>((signed char *)a_base, da[1], (signed char *)b_base, db[1], (signed char *)blocks_info);
}

void loop__2_cuda_(long *cur_dev, long *loop_ref, long a[], long b[])
{
    long da[5], db[5];
    dim3 blocks, threads;
    cudaStream_t stream;
    IndexType *blocks_info;
    void *eps_grid;
    int red_count;
    void *a_base, *b_base;
    float eps;
// Get 'natural' bases
    a_base = dvmh_get_natural_base_(cur_dev, a);
    b_base = dvmh_get_natural_base_(cur_dev, b);
// Fill 'device' headers
    dvmh_fill_header_(cur_dev, a_base, a, da);
    dvmh_fill_header_(cur_dev, b_base, b, db);
// Register reduction for CUDA-execution
    loop_register_red_cuda(loop_ref, 0, &eps_grid, 0);
    loop_red_init(loop_ref, 0, &eps, 0);
// GPU execution
    while (loop_do_cuda(loop_ref, &blocks, &threads, &stream, &blocks_info) != 0) {
        red_count = 1;
        while (red_count * 2 < threads.x * threads.y * threads.z)
            red_count *= 2;
        loop__2<<<blocks, threads, threads.x * threads.y * threads.z * sizeof(float), stream>>>((signed char *)a_base, da[1], (signed char *)b_base, db[1], eps, (signed char *)eps_grid, (signed char *)blocks_info, red_count);
    }
}

void loop__3_cuda_(long *cur_dev, long *loop_ref, long a[], long b[])
{
    long da[5], db[5];
    dim3 blocks, threads;
    cudaStream_t stream;
    IndexType *blocks_info;
    void *a_base, *b_base;
// Get 'natural' bases
    a_base = dvmh_get_natural_base_(cur_dev, a);
    b_base = dvmh_get_natural_base_(cur_dev, b);
// Fill 'device' headers
    dvmh_fill_header_(cur_dev, a_base, a, da);
    dvmh_fill_header_(cur_dev, b_base, b, db);
// GPU execution
    while (loop_do_cuda(loop_ref, &blocks, &threads, &stream, &blocks_info) != 0)
        loop__3<<<blocks, threads, 0, stream>>>((signed char *)a_base, da[1], (signed char *)b_base, db[1], (signed char *)blocks_info);
}

}
